#include "cl_test.h"

void CL_Test::fill_array(float* pfData, int iSize)
{
    int i;
    const float fScale = 1.0f / (float)RAND_MAX;
    for (i = 0; i < iSize; ++i)
    {
        //pfData[i] = fScale * rand();
    	pfData[i] = rand()%10;
    }
}


size_t CL_Test::round_up(int group_size, int global_size)
{
    int r = global_size % group_size;
    if(r == 0)
    {
        return global_size;
    } else
    {
        return global_size + group_size - r;
    }
}

void CL_Test::check_cl_error(cl_uint ret_code){
	if (ret_code!= CL_SUCCESS){
		LOGE("cl error");
	}
}

CL_Test::CL_Test(){
	//iNumElements = 11444777;
	iNumElements = 1024;


    // set and log Global and Local work size dimensions
    szLocalWorkSize = 256;
    szGlobalWorkSize = round_up((int)szLocalWorkSize, iNumElements);  // rounded up to the nearest multiple of the LocalWorkSize
    LOGI("Global Work Size \t\t= %u\nLocal Work Size \t\t= %u\n# of Work Groups \t\t= %u\n\n",
           szGlobalWorkSize, szLocalWorkSize, (szGlobalWorkSize % szLocalWorkSize + szGlobalWorkSize/szLocalWorkSize));

    // Allocate and initialize host arrays
    LOGI( "Allocate and Init Host Mem...\n");
    srcA = new cl_float[szGlobalWorkSize];
    srcB = new cl_float[szGlobalWorkSize];
    dst = new cl_float[szGlobalWorkSize];
    /*srcA = (void *)malloc(sizeof(cl_float) * szGlobalWorkSize);
    srcB = (void *)malloc(sizeof(cl_float) * szGlobalWorkSize);
    dst = (void *)malloc(sizeof(cl_float) * szGlobalWorkSize);
    */
    fill_array((float*)srcA, iNumElements);
    fill_array((float*)srcB, iNumElements);
    for (int i=0;i<szGlobalWorkSize;i++){
    	dst[i] = -1;
    }

	LOGI("cl_test starts\n");

	//Get an OpenCL platform
	check_cl_error(clGetPlatformIDs(1, &cpPlatform, NULL));

	//Get the devices
	check_cl_error(clGetDeviceIDs(cpPlatform, CL_DEVICE_TYPE_GPU, 1, &cdDevice, NULL));

	//Create the context
	cl_int err_code;
	cxGPUContext = clCreateContext(0, 1, &cdDevice, NULL, NULL, &err_code);
	check_cl_error(err_code);

	// Create a command-queue
	cqCommandQueue = clCreateCommandQueue(cxGPUContext, cdDevice, 0, &err_code);
	check_cl_error(err_code);

    // Allocate the OpenCL buffer memory objects for source and result on the device GMEM
    cmDevSrcA = clCreateBuffer(cxGPUContext, CL_MEM_READ_ONLY, sizeof(cl_float) * szGlobalWorkSize, NULL, &err_code);
	check_cl_error(err_code);
    cmDevSrcB = clCreateBuffer(cxGPUContext, CL_MEM_READ_ONLY, sizeof(cl_float) * szGlobalWorkSize, NULL, &err_code);
	check_cl_error(err_code);
    cmDevDst = clCreateBuffer(cxGPUContext, CL_MEM_WRITE_ONLY, sizeof(cl_float) * szGlobalWorkSize, NULL, &err_code);
	check_cl_error(err_code);


    // Read the OpenCL kernel in from source file
	const char* cSourceCL =
			"__kernel void VectorAdd(__global const float* a, __global const float* b, __global float* c, int iNumElements)"
			"{"
			    // get index into global data array
			"    int iGID = get_global_id(0);"

			    // bound check (equivalent to the limit on a 'for' loop for standard/serial C code
			"    if (iGID >= iNumElements)"
			"    {"
			"        return;"
			"    }"

			    // add the vector elements
			"    c[iGID] = a[iGID] + b[iGID];"
			"}";
	szKernelLength = strlen(cSourceCL);
    // Create the program
    cpProgram = clCreateProgramWithSource(cxGPUContext, 1, (const char **)&cSourceCL, &szKernelLength, &err_code);
	check_cl_error(err_code);

	// Build the program
	check_cl_error(clBuildProgram(cpProgram, 0, NULL, NULL, NULL, NULL));

	// Create the kernel
	ckKernel = clCreateKernel(cpProgram, "VectorAdd", &err_code);
	check_cl_error(err_code);

    // Set the Argument values
	check_cl_error(clSetKernelArg(ckKernel, 0, sizeof(cl_mem), (void*)&cmDevSrcA));
	check_cl_error(clSetKernelArg(ckKernel, 1, sizeof(cl_mem), (void*)&cmDevSrcB));
	check_cl_error(clSetKernelArg(ckKernel, 2, sizeof(cl_mem), (void*)&cmDevDst));
	check_cl_error(clSetKernelArg(ckKernel, 3, sizeof(cl_int), (void*)&iNumElements));

	{
		// --------------------------------------------------------
	    // Start Core sequence... copy input data to GPU, compute, copy results back

	    // Asynchronous write of data to GPU device
		check_cl_error(clEnqueueWriteBuffer(cqCommandQueue, cmDevSrcA, CL_FALSE, 0, sizeof(cl_float) * szGlobalWorkSize, srcA, 0, NULL, NULL));
		check_cl_error(clEnqueueWriteBuffer(cqCommandQueue, cmDevSrcB, CL_FALSE, 0, sizeof(cl_float) * szGlobalWorkSize, srcB, 0, NULL, NULL));
	    LOGI("clEnqueueWriteBuffer (SrcA and SrcB)...\n");

	    // Launch kernel
	    check_cl_error(clEnqueueNDRangeKernel(cqCommandQueue, ckKernel, 1, NULL, &szGlobalWorkSize, &szLocalWorkSize, 0, NULL, NULL));
	    LOGI("clEnqueueNDRangeKernel (VectorAdd)...\n");

	    // Synchronous/blocking read of results, and check accumulated errors
	    check_cl_error(clEnqueueReadBuffer(cqCommandQueue, cmDevDst, CL_TRUE, 0, sizeof(cl_float) * szGlobalWorkSize, dst, 0, NULL, NULL));
	    LOGI("clEnqueueReadBuffer (Dst)...\n\n");
	    //--------------------------------------------------------

	    // Compute and compare results for golden-host and report errors and pass/fail
	    /*
	    LOGI("Comparing against Host/C++ computation...\n\n");
	    VectorAddHost ((const float*)srcA, (const float*)srcB, (float*)Golden, iNumElements);
	    bool bMatch = shrComparefet((const float*)Golden, (const float*)dst, (unsigned int)iNumElements, 0.0f, 0);
	    shrLog("%s\n\n", (bMatch == shrTRUE) ? "PASSED" : "FAILED");
	    */

	}
    // Cleanup and leave
	Cleanup ();

	LOGI("cl_test ends\n");
}

void CL_Test::Cleanup ()
{
    // Cleanup allocated objects
    LOGI("Starting Cleanup...\n\n");
	if(ckKernel)clReleaseKernel(ckKernel);
    if(cpProgram)clReleaseProgram(cpProgram);
    if(cqCommandQueue)clReleaseCommandQueue(cqCommandQueue);
    if(cxGPUContext)clReleaseContext(cxGPUContext);
    if(cmDevSrcA)clReleaseMemObject(cmDevSrcA);
    if(cmDevSrcB)clReleaseMemObject(cmDevSrcB);
    if(cmDevDst)clReleaseMemObject(cmDevDst);

    // Free host memory
    /*
    free(srcA);
    free(srcB);
    free (dst);
    */
    delete[] srcA;
    delete[] srcB;
    delete[] dst;
}
